Stephen Jones, GTC 2022
软件开发的第一个时代是单线程时代(持续至约2007年),其特点是顺序执行代码(Straight-Line Code)。
在这个时代,程序指令按顺序一步一步执行。即便是在现代,许多复杂算法的逻辑流程也可以被视为顺序执行代码。例如,一个典型的深度学习模型(如下图所示的Transformer架构)就包含了一系列顺序执行的层,如输入嵌入、多头注意力、前馈网络等。
下图展示了一个高密度计算节点的内部结构,这通常是现代数据中心和AI超级计算机的基本构建单元。图中可以清晰地看到多个GPU模块、高速互联组件以及先进的散热系统(如铜制散热片和热管)。这种设计旨在将强大的计算能力集成在单个紧凑的服务器单元中。
计算需求已从单个节点扩展到整个数据中心。通过将前述的高密度计算节点集成到机架中,并部署大量机架,可以构建出用于处理大规模AI训练和高性能计算(HPC)工作负载的数据中心级计算集群。
在计算架构中,管理数据局部性(Locality)是一个由来已久的核心问题。下图示意了这一基本概念:处理器(方块)访问与其物理位置更近的内存(矩形)时,速度更快、延迟更低。图中高亮的蓝色部分展示了一组处理器及其紧密耦合的本地内存。在扩展到数据中心规模时,如何高效地管理跨节点、跨机架的数据局部性,对于系统整体性能至关重要。
在过去的十年里,GPU 的并行处理能力实现了巨大的飞跃。以 NVIDIA Hopper H100 GPU 为例,其完整芯片包含 132 个流式多处理器(SMs),而十年前的 Kepler GK110 完整芯片仅有 15 个 SMs。这意味着,整个 GK110 芯片的计算核心数量,如今仅相当于 H100 芯片的一小部分。这种规模的增长要求编程模型也随之演进,以便开发者能够有效地利用如此大规模的并行硬件。
CUDA 编程模型采用层次化的结构来组织和管理并行任务,其核心概念是 Grid(网格)、Block(线程块) 和 Thread(线程)。
一个完整的计算任务,在 CUDA 中被称为一个 Grid。可以将其想象成需要处理的整个工作负载,例如处理一张完整的图像。
为了实现并行处理,整个 Grid 被划分为一个由多个大小相等的 Blocks 组成的网格。每个 Block 负责处理整个任务的一部分。
每个 Block 内部由许多个 Threads 组成。一个 Block 可以看作一个独立的程序,其内部的多个线程协同执行。这种模型允许将大规模问题分解为可以在 GPU 的众多核心上并行执行的小任务。
为了更好地利用现代 GPU 架构的物理局部性,Hopper 架构引入了一个新的层次:线程块集群 (Thread Block Cluster)。
一个线程块集群是一组被共同调度到相邻多处理器上的线程块的集合。这在原有的 Grid → Block → Thread 层次结构中增加了一个新的层级,变为 Grid → Cluster → Block → Thread。
线程块集群的设计初衷是为了充分利用 GPU 硬件的物理局部性。通过将一个集群内的所有线程块调度到物理上相邻的 SMs 上,可以实现以下优势:
- 保证协同定位的线程块 (Guaranteed co-located blocks):集群内的块在物理上彼此靠近。
- 新的保证并发层级 (New tier of guaranteed concurrency):集群内的所有块保证同时运行。
- 快速数据交换与同步 (Fast data exchange & sync):由于物理上的邻近性,集群内的线程块可以更高效地进行数据共享和同步。
在程序中构建和使用线程块集群需要遵循特定的编程规范。
__cluster_dims__(x, [y, [z]]) 来为一个 CUDA 核函数 (kernel) 指定其所需的集群维度。4x2x1(即 8个)线程块组成的集群:```c++
// 定义一个 4x2x1 的8块集群
cluster_dims(4, 2, 1)
global void hellocluster()
{
// 获取当前集群的 cooperative group
cooperative_groups::cluster_group cluster = this_cluster();
// 在集群内进行同步
cluster.sync();
printf("Hello from cluster elem %d\n", cluster.cluster_rank());
}
```
- 启动配置:此外,还引入了新的可扩展启动 API,允许在运行时配置集群参数。
Tensor Memory Accelerator (TMA) 是一种用于异步数据移动的硬件加速单元。它支持硬件加速的双向批量拷贝,可以在全局内存(Global Memory)与共享内存(Shared Memory)之间,以及在集群内分布式共享内存(DSMEM)之间进行数据传输。TMA 使用异步事务屏障(Asynchronous Transaction Barrier)来跟踪操作的完成情况。
TMA 的一个关键应用是实现低延迟的单边数据传输。它支持在集群内部进行单边批量数据拷贝,并在拷贝到全局内存时支持元素级的规约(reduction)操作。
一个典型的例子是在一个集群内进行单边光环交换(halo exchange)。如下图所示,在一个2x2的线程块集群(Cluster of Blocks)中,TMA可以通过自同步事务(self-synchronizing transactions)在集群内的不同线程块之间高效交换边界数据(halo区域),其速度比传统方法快7倍。
TMA 带来了更快、更灵活的异步拷贝功能,它作为对现有 cuda::memcpy_async() 函数的直接增强。其核心优势包括:
- 单线程触发:单个线程即可触发任意大小的数据拷贝,无需编写循环或使用协作式拷贝(collective copying)。
- 聚合屏障:多个拷贝操作可以汇集到一个集群范围的异步事务屏障上。
如下图所示,TMA 单元在 H100 SM 内部,负责处理从 GPU 内存到共享内存的数据传输。通过一个新的异步事务屏障机制(Arrive -> Wait),可以有效跟踪数据传输的完成状态。编程模型上,开发者可以使用带有屏障对象的新版 cuda::memcpy_async 语法来显式启用 TMA 的高级功能,而旧的语法在 H100 上也会自动利用 TMA 进行增强。
TMA 能够硬件加速多维张量的拷贝,最高支持5阶张量(rank 5)。其特性包括:
- 自动地址生成:自动处理多维张量的步长(stride)和地址计算。
- 边界填充:为越界访问提供边界填充(Boundary padding)。
- “发射后不管”:仅需单个线程发起拷贝指令,后续所有复杂操作均由 TMA 硬件处理,开发者无需编写迭代或边界检查代码。
下图展示了 TMA 如何从一个大的多维张量中拷贝一个子区域(sub-region)。TMA 会自动处理张量访问步长、块宽度/高度以及必要的自动填充。
Cooperative Groups 是一个用于显式表达线程间协作和同步的编程模型。它将GPU的执行层级结构暴露给开发者,允许在不同粒度上进行编程。
GPU 的自然执行层级从粗到细依次为:
- Grid of work: 整个内核的工作网格。
- Cluster of Blocks: 线程块的集群。
- Block of Threads: 线程块。
- Warp of Threads: 线程束。
开发者可以通过相应的API获取每个层级的句柄,例如 this_grid()、this_cluster()、this_thread_block() 和 tiled_partition<32>()。
Cooperative Groups 在不同的执行层级上支持不同的协作操作:
- Grid 和 Cluster 层级: 主要支持同步(Synchronization)。需要注意的是,要使用 grid 范围的协作操作,Grid 必须以协作方式启动(cooperatively launched)。
- Block 和 Warp 层级: 支持更丰富的操作,包括同步(Synchronization)、规约(Reduction)和前缀和(Prefix-sum)。
- Warp 及更小粒度的组: 除了上述操作外,还支持 shfl, any, all, ballot, match 等 warp 内置函数。
该模型还支持由多个 warp 组成的组进行协作操作。支持的 thread_block_tile 尺寸包括64、128、256和512。例如,可以通过 cooperative_groups::thread_block_tile<128> 来创建一个由128个线程(即4个warp,一个 quadwarp)组成的组,并在此组上执行协作操作。
Cooperative Groups 通过C++的类型系统,可以创建类型安全的并行函数库。函数可以明确声明其期望的线程数量。例如,一个为128个线程设计的基数-128 FFT(Radix-128 FFT)函数,可以将其接口参数类型定义为 const thread_block_tile<128>&。这样,编译器就能确保该函数只能被一个包含128个线程的 thread_block_tile 对象调用,从而实现了编译时的安全性,避免了运行时错误。
// 内核函数将线程块划分为128线程的tile
__global__ void kernel(...) {
thread_block_tile<128> tile128 = tiled_partition<128>(this_thread_block());
fft128(tile128, ...); // 调用FFT函数
}
// 一个必须由128个线程调用的设备函数
__device__ void fft128(const thread_block_tile<128>& group, ...) {
do_fft(...);
group.sync(); // 同步组内所有128个线程
}
Cooperative Groups 可用于实现复杂的分层计算模式。以下是一个流体模拟的例子:
1. 启动一个由 4x4 个 cluster 组成的协作式 Grid,每个 cluster 包含 4x2 个 block。
2. 使用协作式 cluster 索引,从邻近的 cluster 拉取数据。
3. 在每个 block 内部进行局部求解。
4. 在 cluster 层面进行同步,并利用分布式共享内存(distributed shared memory)进行块间通信和求解。
5. 在 grid 层面进行同步和求解,然后返回第2步,开始下一次迭代。
生产者/消费者(Producer/Consumer)模型可以在 Grid、Cluster 或 Block 等任意层级上实现。
- Grid 级别: 将整个 GPU 上的 Clusters 划分为 "生产者" 和 "消费者" 角色。它们通过全局内存交换数据,并使用 this_grid().sync() 进行同步。
- Cluster 级别: 在一个 Cluster 内部,将 Blocks 划分为 "生产者" 和 "消费者" 角色。它们通过分布式共享内存或全局内存交换数据,并使用 this_cluster().sync() 进行同步。
- Block 级别: 在一个 Block 内部,将 Warps 划分为 "生产者" 和 "消费者" 角色。它们通过共享内存或全局内存交换数据,并使用 this_thread_block().sync() 进行同步。
该模式在 Transformer 模型的计算中也有应用,例如可以将模型中的一部分层指定为生产者,另一部分为消费者。
这是一个将生产者/消费者任务并行应用于量化金融领域的实例。
- Longstaff Schwartz 模型: 一种通过蒙特卡洛模拟和反向迭代来确定金融期权价值的定价技术。
- 分析: 分布式共享内存(Distributed shared memory)的更大容量为高效内核的编程带来了新方法。线程块级别的生产者/消费者模型现在可以在共享内存中实现,并利用快速的 cluster 同步。硬件加速的 cluster 同步机制易于使用且高效。
- 性能: 实验表明,在 H100 平台上,与不使用 cluster 的方案相比,使用 cluster 同步(Cluster.sync())可以将 Longstaff Schwartz 模型的吞吐量提升 2.7倍。如果进一步使用 cluster 和异步屏障(Barrier),性能还能再提升 10%。
CUDA 11.7 在整个工具链和库中增加了对 __int128 类型的支持。
- 兼容性: 可与兼容的主机编译器(如 gcc, clang, icc)一起使用,或通过运行时编译(NVRTC)使用。
- 完整支持: 支持算术、逻辑、位运算、数学运算、库函数和开发工具等。
开发者可以在主机代码和设备代码(kernel)中无缝使用128位整数。
为了缩短应用的构建时间,NVRTC 和 PTX JIT(即时编译)的锁机制从全局锁(global locking)改进为分阶段锁(per-stage locking)。这一改进使得在使用多个CPU线程进行编译时,能够实现并发编译。
- 效果: 从顺序编译流程转变为流水线式(Pipelined)编译流程,显著提升了编译速度。
- 性能: 测试表明,在 CUDA 11.5 中,使用4个CPU线程进行编译的速度比使用单个线程快一倍以上(从44秒缩短至20秒)。
__grid_constant__ 注解,用于内核参数。该注解允许编译器避免为只读参数创建每个线程的私有副本,从而优化资源使用。#pragma 指令(如 nv_diag_suppress, nv_diag_warning 等)来在代码中动态地控制编译器的诊断信息(错误、警告)。-arch=all: 为所有支持的架构生成代码。-arch=all-major: 为所有支持的主要架构版本生成代码。-arch=native: 为系统上所有可见的GPU生成代码。即时编译链接(JIT Linking)与链接时优化(LTO)相结合,能够在运行时将对象文件与库进行组合和优化。该技术允许在a.out(可执行文件)与cuFFTDx内核库之间进行调用,最终生成一个完全优化的内核。
性能对比:
在A100(40GB)GPU上,JIT-LTO回调相较于间接回调(Indirect Callback)在不同FFT规模下展现了显著的性能提升,最高可达1.92倍。
nvcc -dlto 编译 .cu 文件生成可重定位的对象文件,运行时JIT LTO将此对象文件与cuFFTDx内核库链接并进行联合优化,生成一个高效的、完全优化的内核。
设备端代码直接调用数学库(如cuFFTDx)可以显著加速小规模问题的处理。
性能对比:设备端调用 vs. 主机端启动
在A100(80GB)GPU上,与传统的主机端启动cuFFT相比,从设备端代码调用的cuFFTDx在处理小尺寸FFT时性能优势尤为明显,TFlop/s性能提升显著。
支持设备端调用的接口库:
通过在高性能数学库中调用自定义内核代码,可以实现更复杂的计算流程。
cuFFTDx: 快速傅里叶变换。cuBLASDx: GPU加速的基础线性代数子程序库 (BLAS-1, 2, 3)。cuSOLVERDx: GPU加速的稠密线性代数库。应用案例:带内联卷积的FFT
在带内联卷积的FFT计算中,cuFFTDx的性能优于标准的cuFFT以及使用回调机制的cuFFT。
CUTLASS 是一个用于深度学习和高性能计算的 CUDA C++ 模板库。它提供了在各种范围和规模下最优化的 CUDA C++ 矩阵运算模板。
CUTLASS 在不同抽象层次上提供了 Tensor Core 加速的操作。
- Device: GEMM、卷积、归约等,支持所有数据类型、SIMT、Tensor Core及所有架构。
- Kernel: GEMM、批处理GEMM、卷积、归约、融合操作等。
- Thread Block: 流水线矩阵乘法、Epilogue、到张量的集合操作、卷积矩阵访问。
- Warp: 张量核乘加操作、高效访问置换内存布局。
- Thread: 数值转换、<functional> 操作符、复数<T>、快速数学算法。
- Architecture Intrinsic: 包装了架构特定的PTX指令(如 mma, cp.async, ldmatrix, cvt)。
CUTLASS 采用一个三步TF32序列,在A100 GPU上实现了48 TFLOPs的FP32精度计算性能。
性能与精度分析:
- 性能: 3步TF32 GEMM的性能远超A100的FP32峰值性能(19.5 TFLOPs)。
- 精度: 相较于标准的FP32计算,3xTF32方法的相对误差略高,并随问题规模(GEMM-K维度)增大而增加。
CUDA生态系统提供了一系列丰富的开发工具,涵盖调试、性能分析、代码修正和IDE集成。
新的网络性能分析功能可以拦截并追踪对UCX(Unified Communication X)协议层的调用。这使得开发者能够深入了解MPI通信背后的底层UCP(Unified Communication Protocol)操作。
Nsight Systems现在可以直接展示网卡(NIC)的性能指标,帮助分析多节点应用的通信性能。
$ nsys profile --nic-metrics=[true|false] ... 来控制是否采集网卡指标。
Nsight Compute 提供了深入分析CUDA内核性能的强大工具。
NVTX (NVIDIA Tools Extension) 是一个用于在代码中插入标记以供工具进行分析的注解库。
push/pop 对记录一个时间段。范围形成一个栈结构,可以展示嵌套的细节。#include <nvtx3/nvToolsExt.h>
void LaunchKernel() {
nvtxRangePush(__FUNCTION__);
Kernel<<<20000, 256>>>();
nvtxRangePop();
}
void Example() {
nvtxRangePush("Creating CUDA context");
cudaFree(0);
nvtxRangePop();
nvtxRangePush("Launching & waiting for kernels");
for (int i = 0; i < 3; ++i) {
LaunchKernel();
}
cudaDeviceSynchronize();
nvtxRangePop();
}
NVTX的注解代码会直接映射到性能分析工具的时间线视图中,形成层次化的事件范围,极大地增强了代码执行流程的可读性。
nvtxRangePush("Creating CUDA context"); 在时间线上生成一个名为 "Creating CUDA context" 的范围。nvtxRangePush("Launching & waiting for kernels"); 创建了一个父范围。LaunchKernel 中调用的 nvtxRangePush(__FUNCTION__); 则在父范围内部生成了多个嵌套的 "LaunchKernel" 子范围。
一个简单的CUDA "Hello world"程序经过nvcc编译后,会生成一个包含主机和设备代码的可执行文件。
// Hello world example
__global__ void kernel() {
printf("Hello, CUDA\n");
}
void main() {
kernel<<<1, 1>>>();
cudaDeviceSynchronize();
}
编译命令:nvcc example.cu -o example
与传统的CPU程序不同,编译生成的CUDA可执行文件包含多个部分:
- CPU代码区 (CPU code section): 包含主机端代码,如main()函数。
- GPU代码区 (GPU code section): 包含编译为特定GPU架构机器码的设备端代码,如kernel()。
- PTX JIT区 (PTX JIT section): 包含PTX(Parallel Thread Execution)中间代码。这使得程序可以在运行时即时编译(JIT),以兼容未来或编译时未指定的GPU架构。
当链接大型库(如 cuDNN)时,该库会向最终的二进制文件中添加大量的内核 (kernels) 和 PTX 文件,导致二进制文件体积显著增大。
nvcc example.cu -o example -lcudnn
当一个CUDA应用程序启动时,其标准加载流程如下:
1. 从磁盘加载到主机内存:操作系统将整个二进制文件从磁盘加载到CPU主机内存中。由于所有代码段都被完整加载,这会导致较大的主机内存占用(例如,1,866MB)。
2. 从主机内存上传到设备内存:CUDA 驱动程序在初始化阶段,会将二进制文件中包含的所有 GPU 内核从主机内存上传到 GPU 设备内存中。这个过程可能非常耗时(例如,1.6秒),并且会占用大量的设备内存(例如,1,245MB),即使应用程序在整个生命周期中可能只使用到其中的一小部分内核。
为了优化上述标准加载流程带来的高昂开销,引入了延迟加载(Lazy Loading)机制。其核心思想是:仅在函数(内核)首次被引用或访问时,才将其从主机内存上传到设备内存。
采用延迟加载可以带来显著的优化效果:
- 减少二进制加载时间:由于初始阶段无需上传所有内核,程序启动速度加快。
- 降低设备内存占用:GPU 内存中只保留当前活跃的内核,大幅减少了静态内存占用。
- 降低主机内存占用:同样可以观察到主机内存占用的减少。
延迟加载功能可以通过设置一个环境变量来激活:
set CUDA_MODULE_LOADING=LAZY
在一个简单的 cuDNN 卷积测试中(环境:Ubuntu 20.04, A100 40GB),延迟加载与标准加载的性能对比如下:
- 二进制加载时间 (s):标准加载为 1.6秒,延迟加载为 0.8秒,时间减少 50%。
- 设备内存占用 (MB):标准加载为 1,245MB,延迟加载为 435MB,占用减少约 65%。
- 主机内存占用 (MB):标准加载为 1,866MB,延迟加载为 1,229MB,占用减少约 34%。
实验数据明确显示,延迟加载能够有效优化应用程序的启动时间和内存资源利用率。
CUDA Graphs 提供了对复杂工作流进行高效管理和执行的能力。新的特性进一步增强了其灵活性和控制力。
该功能允许用户在图启动(launch)之前动态修改执行流程,通过禁用图中的某些节点来实现。
该功能允许为任务图中的特定分支或路径创建高优先级,从而影响调度和执行顺序。